<!DOCTYPE html>
<!--[if IE 8]><html class="no-js lt-ie9" lang="en" > <![endif]-->
<!--[if gt IE 8]><!--> <html class="no-js" lang="en" > <!--<![endif]-->
<head>
  <meta charset="utf-8">
  <meta http-equiv="X-UA-Compatible" content="IE=edge">
  <meta name="viewport" content="width=device-width, initial-scale=1.0">
  
  <meta name="author" content="PG-Strom Development Team">
  <link rel="shortcut icon" href="../img/favicon.ico">
  <title>Python連携 - PG-Strom Manual</title>
  <link href='https://fonts.googleapis.com/css?family=Lato:400,700|Roboto+Slab:400,700|Inconsolata:400,700' rel='stylesheet' type='text/css'>

  <link rel="stylesheet" href="../css/theme.css" type="text/css" />
  <link rel="stylesheet" href="../css/theme_extra.css" type="text/css" />
  <link rel="stylesheet" href="//cdnjs.cloudflare.com/ajax/libs/highlight.js/9.12.0/styles/github.min.css">
  <link href="//fonts.googleapis.com/earlyaccess/notosansjp.css" rel="stylesheet">
  <link href="//fonts.googleapis.com/css?family=Open+Sans:600,800" rel="stylesheet">
  <link href="../custom.css" rel="stylesheet">
  
  <script>
    // Current page data
    var mkdocs_page_name = "Python\u9023\u643a";
    var mkdocs_page_input_path = "python.md";
    var mkdocs_page_url = null;
  </script>
  
  <script src="../js/jquery-2.1.1.min.js" defer></script>
  <script src="../js/modernizr-2.8.3.min.js" defer></script>
  <script src="//cdnjs.cloudflare.com/ajax/libs/highlight.js/9.12.0/highlight.min.js"></script>
  <script>hljs.initHighlightingOnLoad();</script> 
  
</head>

<body class="wy-body-for-nav" role="document">

  <div class="wy-grid-for-nav">

    
    <nav data-toggle="wy-nav-shift" class="wy-nav-side stickynav">
      <div class="wy-side-nav-search">
        <a href=".." class="icon icon-home"> PG-Strom Manual</a>
        <div role="search">
  <form id ="rtd-search-form" class="wy-form" action="../search.html" method="get">
    <input type="text" name="q" placeholder="Search docs" />
  </form>
  [<strong>Japanese</strong> | <a href="../../python/"    style="color: #cccccc">English</a>]
</div>
      </div>

      <div class="wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="main navigation">
	<ul class="current">
	  
          
            <li class="toctree-l1">
		
    <a class="" href="..">はじめに</a>
	    </li>
          
            <li class="toctree-l1">
		
    <a class="" href="../install/">インストール</a>
	    </li>
          
            <li class="toctree-l1">
		
    <span class="caption-text">利用ガイド</span>
    <ul class="subnav">
                <li class="">
                    
    <a class="" href="../operations/">基本的な操作</a>
                </li>
                <li class="">
                    
    <a class="" href="../sys_admin/">システム管理</a>
                </li>
                <li class="">
                    
    <a class="" href="../brin/">BRINインデックス</a>
                </li>
                <li class="">
                    
    <a class="" href="../partition/">パーティション</a>
                </li>
                <li class="">
                    
    <a class="" href="../postgis/">PostGIS対応</a>
                </li>
                <li class="">
                    
    <a class="" href="../troubles/">トラブルシューティング</a>
                </li>
    </ul>
	    </li>
          
            <li class="toctree-l1">
		
    <span class="caption-text">先進機能</span>
    <ul class="subnav">
                <li class="">
                    
    <a class="" href="../ssd2gpu/">GPUダイレクトSQL</a>
                </li>
                <li class="">
                    
    <a class="" href="../arrow_fdw/">Apache Arrow</a>
                </li>
                <li class="">
                    
    <a class="" href="../gstore_fdw/">GPUメモリストア</a>
                </li>
                <li class=" current">
                    
    <a class="current" href="./">Python連携</a>
    <ul class="subnav">
            
    <li class="toctree-l3"><a href="#_1">背景</a></li>
    

    <li class="toctree-l3"><a href="#cupy">cuPyを用いたデータ交換</a></li>
    
        <ul>
        
            <li><a class="toctree-l4" href="#cupy_strom">cupy_stromのインストール</a></li>
        
            <li><a class="toctree-l4" href="#cupygpu">cuPyを用いたカスタムGPUカーネルの実行</a></li>
        
            <li><a class="toctree-l4" href="#plpython">PL/Pythonユーザ定義関数からの利用</a></li>
        
        </ul>
    

    </ul>
                </li>
    </ul>
	    </li>
          
            <li class="toctree-l1">
		
    <span class="caption-text">リファレンス</span>
    <ul class="subnav">
                <li class="">
                    
    <a class="" href="../ref_types/">データ型</a>
                </li>
                <li class="">
                    
    <a class="" href="../ref_devfuncs/">関数と演算子</a>
                </li>
                <li class="">
                    
    <a class="" href="../ref_sqlfuncs/">SQLオブジェクト</a>
                </li>
                <li class="">
                    
    <a class="" href="../ref_params/">GUCパラメータ</a>
                </li>
    </ul>
	    </li>
          
            <li class="toctree-l1">
		
    <a class="" href="../release_note/">リリースノート</a>
	    </li>
          
        </ul>
      </div>
      &nbsp;
    </nav>

    <section data-toggle="wy-nav-shift" class="wy-nav-content-wrap">

      
      <nav class="wy-nav-top" role="navigation" aria-label="top navigation">
        <i data-toggle="wy-nav-top" class="fa fa-bars"></i>
        <a href="..">PG-Strom Manual</a>
      </nav>

      
      <div class="wy-nav-content">
        <div class="rst-content">
          <div role="navigation" aria-label="breadcrumbs navigation">
  <ul class="wy-breadcrumbs">
    <li><a href="..">Docs</a> &raquo;</li>
    
      
        
          <li>先進機能 &raquo;</li>
        
      
    
    <li>Python連携</li>
    <li class="wy-breadcrumbs-aside">
      
    </li>
  </ul>
  <hr/>
</div>
          <div role="main">
            <div class="section">
              
                <p>本章では、共通データフレームとPythonスクリプト（PL/Pythonを含む）を用いて、機械学習や統計解析エンジンとPostgreSQLを接続する、あるいはこれらのエンジンを In-database で実行する方法について説明します。</p>
<h1 id="_1">背景</h1>
<p>過去のバージョンのPG-Stromにおいては、データベースと機械学習・統計解析エンジンの連携のためにPL/CUDAという仕組みを提供していました。これは、SQLのユーザ定義関数としてCUDA Cのコードを記述し、計算集中的なワークロードをGPUの持つ数千コアのプロセッサで処理する事を可能にするもので、場合によってはCPUの百倍以上の処理速度を実現する事もありました。</p>
<p>しかし一方で、機械学習や統計解析の領域では多くの場合Python言語から利用できるモジュールを組み合わせてアプリケーションを構築するのが一般的な流れとなり、個々のデータサイエンティストがCUDA C言語で並列プログラムを記述するというケースは稀でした。</p>
<p>PG-Strom v2.3以降では、データベースの内容をPython向け機械学習・統計解析用モジュールと共通形式のデータフレームとして交換する事が可能で、データをエクスポートする事なくGPUの計算能力を利用するというPL/CUDAの特徴を引き継ぎつつ、Python言語による高い生産性を両立しています。</p>
<p><img alt="Data Processing Lifecycle" src="../img/data-processing-lifecycle.png" /></p>
<p>本章で説明する機能は、基本的にはオンメモリ処理が可能なデータ量を対象としたものです。
したがって、テラバイトを越えるような大きさのデータをPythonから参照可能なデータフレームとして見せるために使用すべきではありません。</p>
<p>これは例えば、大規模な生ログデータはSSD-to-GPU Direct SQLを用いて前処理し、次にこれを共通データフレームを介してPython上の機械学習エンジンに入力するといった使い方を想定しています。</p>
<h1 id="cupy">cuPyを用いたデータ交換</h1>
<p><a href="https://cupy.chainer.org/">cuPy</a>とは、Python環境において行列演算ライブラリとして広く使われているnumPyと共通のAPI群をもち、その演算をGPU上で実行するためのライブラリです。</p>
<p>cuPyが計算に利用する行列データはGPU上に配置され、<code>cupy.ndarray</code>型のオブジェクトとしてスクリプト上から操作する事が可能です。cuPyのインストールやAPIリファレンスなどは、<a href="https://docs-cupy.chainer.org/">公式ドキュメント</a>をご覧ください。</p>
<p>PG-StromはArrow_Fdw外部テーブルに保持されたデータをGPU上に展開し、Pythonスクリプトの実行環境に当該メモリ領域をマップする事ができます。これにより、データベースからデータをエクスポートする事なくスマートにデータ交換を行う事ができます。</p>
<p>PG-StromのSQL関数<code>pgstrom.arrow_fdw_export_cupy()</code>はGPUデバイスを割り当て、Arrow_Fdw外部テーブルの内容をロードします。この関数は外部テーブル、ロードすべき列名、およびターゲットのGPUデバイスID（省略可）を引数に取り、当該GPUバッファの識別子を<code>text</code>で返します。</p>
<p>Pythonスクリプト側では<code>cupy_strom.ipc_import()</code>にこの識別子を与えて、GPUデバイスメモリ上に展開済みの<code>cupy.ndarray</code>をオープンします。以下に実例を示します。</p>
<p><strong>PostgresSQL側のオペレーション</strong></p>
<pre><code>=# CREATE FOREIGN TABLE ft (
     id  int,
     x   real,
     y   real,
     z   real
   ) SERVER arrow_fdw
     OPTIONS (file '/dev/shm/ftest.arrow', writable 'true');

=# INSERT INTO ft (SELECT x, pgstrom.random_int(0,1,10000)::float/100.0,
                             pgstrom.random_int(0,1,10000)::float/100.0,
                             pgstrom.random_int(0,1,10000)::float/100.0
                     FROM generate_series(1,5) x);

=# SELECT * FROM ft;
 id |   x   |   y   |   z
----+-------+-------+-------
  1 | 51.61 | 73.23 |  7.53
  2 | 49.73 | 29.75 | 37.31
  3 | 61.15 | 55.15 | 36.07
  4 | 23.76 | 40.76 | 51.74
  5 | 61.43 | 86.87 | 47.64
(5 rows)
</code></pre>

<p><strong>Python側のオペレーション</strong></p>
<pre><code>import psycopg2
import cupy
import cupy_strom

conn = psycopg2.connect(&quot;host=localhost dbname=postgres&quot;)
curr = conn.cursor()
curr.execute(&quot;select pgstrom.arrow_fdw_export_cupy('ft','{x,y,z}'::text[])&quot;)
row = curr.fetchone()

X = cupy_strom.ipc_import(row[0])

print(X)

conn.close()
</code></pre>

<p><strong>実行結果</strong></p>
<pre><code>$ python ftest.py
[[51.61 49.73 61.15 23.76 61.43]
 [73.23 29.75 55.15 40.76 86.87]
 [ 7.53 37.31 36.07 51.74 47.64]]
</code></pre>

<p>上記の例は、PythonスクリプトからPostgreSQLへ接続し、SQL関数<code>pgstrom.arrow_fdw_export_cupy</code>を用いて外部テーブル<code>ft</code>の列<code>x</code>、<code>y</code>および<code>z</code>の3列から成る<code>cupy.ndarray</code>用のGPUバッファを作成しています。続いて、その関数の返り値である識別子を<code>cupy_strom.ipc_import</code>関数に渡す事で、Pythonから利用可能な<code>cupy.ndarray</code>オブジェクトを生成しています。</p>
<p>いったん<code>cupy.ndarray</code>オブジェクトが生成された後は、既存の cuPy のAPI群を用いてこのGPUバッファを操作する事ができます。ここでは僅か5行x3列のデータを扱いましたが、これが10億行のデータになったとしても、同様にPostgreSQLとPythonスクリプトの間でデータ交換を行う事ができます。</p>
<p>割り当てたGPUバッファはセッションの終了時に自動的に解放されます。セッション終了後もGPUバッファを保持し続けたい場合は、代わりに<code>pgstrom.arrow_fdw_export_cupy_pinned</code>を使用してGPUバッファを割り当てます。この場合、明示的に<code>pgstrom.arrow_fdw_unpin_gpu_buffer</code>を呼び出してピンニング状態を解除するまでは、GPUデバイスメモリを占有し続ける事に留意してください。</p>
<h2 id="cupy_strom">cupy_stromのインストール</h2>
<p>前述の操作に必要な<code>cupy_strom</code>パッケージは、<code>setup.py</code>スクリプトを用いて以下のようにインストールする事ができます。</p>
<pre><code>$ sudo pip3 install --upgrade numpy cupy cython
$ git clone https://github.com/heterodb/pg-strom.git
$ cd pg-strom/python
$ python3 setup.py 
</code></pre>

<h2 id="cupygpu">cuPyを用いたカスタムGPUカーネルの実行</h2>
<p>cuPyは豊富な行列演算APIを持っており、これらを利用する事でCUDA CのプログラミングなしにGPUの計算能力を引き出す事も可能ですが、ユーザが作成したカスタムのGPUカーネル関数を定義し、実行する事も可能です。</p>
<p>以下の例は<code>cupy.RawKernel</code>を使用してカスタムのGPUカーネルを定義したもので、入力値('X')の各列ごとにその平均値を導出するというものです。
<code>cupy.RawKernel</code>オブジェクトの作成には、CUDA Cで記述されたGPUカーネルのソースコードと、GPUカーネルのエントリポイントとなるデバイス関数名が必要で、これは<code>__call__</code>メソッドの呼び出し時に実行時コンパイルが行われます。（ビルド済みバイナリがキャッシュに見つからなければ）</p>
<p><code>__call__</code>メソッドの引数は順にグリッドの大きさ、ブロックの大きさ、およびGPUカーネル関数への引数です。詳細な説明は省きますが、入力値<code>X</code>を2048個の要素ごとに領域分割し、1024個のスレッドが相互に協調動作を行い、11回のステップで各ブロックの総和を計算します。これが各ブロック毎に並列に実行され、最終的に出力バッファ<code>Y</code>には列ごとの総和が格納される事になります。</p>
<pre><code>import psycopg2
import cupy
import cupy_strom

// connect to PostgreSQL, and get identifier of GPU buffer
conn = psycopg2.connect(&quot;host=localhost dbname=postgres&quot;)
curr = conn.cursor()
curr.execute(&quot;select pgstrom.arrow_fdw_export_cupy('ft','{x,y,z}'::text[])&quot;)
row = curr.fetchone()

// import GPU buffer using the identifier string
X = cupy_strom.ipc_import(row[0])
nattrs = X.shape[0]
nitems = X.shape[1]
gridSz = (nitems + 2047) &gt;&gt; 11;
Y = cupy.zeros((nattrs))

// source code of the custom GPU kernel
source='''
extern &quot;C&quot; __global__
           __launch_bounds__(1024)
void
kern_gpu_sum(double *y, const float *x, int nitems)
{
    __shared__ float lvalues[2048];
    int     gridSz = (nitems + 2047) / 2048;
    int     colIdx = blockIdx.x / gridSz;
    int     rowBase = (blockIdx.x % gridSz) * 2048;
    int     localId = 2 * threadIdx.x;
    int     i, k;

    // Load values to local shared buffer
    x += colIdx * nitems;
    for (i=threadIdx.x; i &lt; 2048; i+=blockDim.x)
        lvalues[i] = (rowBase + i &lt; nitems ? x[rowBase + i] : 0.0);
    __syncthreads();

    // Run reduction operations
    for (k=0; k &lt; 11; k++)
    {
        int     mask = ((1 &lt;&lt; k) - 1);

        if ((threadIdx.x &amp; mask) == 0)
            lvalues[localId] += lvalues[localId + (1&lt;&lt;k)];
        __syncthreads();
    }
    // Write back the total sum
    if (threadIdx.x == 0)
        atomicAdd(&amp;y[colIdx], lvalues[0]);
}
'''
kern = cupy.RawKernel(source, 'kern_gpu_sum')
kern.__call__((gridSz * nattrs,1,1),
              (1024,1,1),
              (Y,X,nitems))
print(Y / nitems)

conn.close()
</code></pre>

<p><strong>実行結果</strong></p>
<pre><code>=# SELECT pgstrom.arrow_fdw_truncate('ft');

=# INSERT INTO ft (SELECT x, pgstrom.random_int(0,1,10000)::float/100.0,
                             pgstrom.random_int(0,-7500,2500)::float/100.0,
                             pgstrom.random_int(0,5000,15000)::float/100.0
                     FROM generate_series(1,1000000) x);

=# SELECT avg(x), avg(y), avg(z) FROM ft;
       avg        |        avg        |       avg
------------------+-------------------+------------------
 50.0225953391276 | -24.9964806686448 | 100.037490822002
(1 row)
</code></pre>

<pre><code>$ python ftest.py
[ 50.02259536 -24.99648063 100.03749086]
</code></pre>

<p>意図的に各列の分布をずらしたテストデータによる平均値の計算ですが、GPUバッファを介してcuPyとデータ交換を行い、カスタムGPUカーネルで計算した平均値と、SQLでの計算結果が一致している事が分かります。</p>
<h2 id="plpython">PL/Pythonユーザ定義関数からの利用</h2>
<p>PostgreSQLではPython言語によるユーザ定義関数の記述が可能で、標準で同梱されている<a href="https://www.postgresql.jp/document/current/html/plpython.html">PL/Python</a>パッケージがその機能を提供します。</p>
<p><code>CREATE FUNCTION</code>構文の<code>LANGUAGE</code>句に<code>plpython3u</code>と指定する事で、そのユーザ定義関数はPythonで記述されている事を示します。</p>
<p>以下にPL/Pythonユーザ定義関数の例を示します。平均値を求めるGPUカーネル関数にはもう一度登場してもらう事にします。</p>
<p>PL/Pythonユーザ定義関数の引数は、適切なPythonデータ型にマッピングされます。ここでは、SQL関数<code>pgstrom.arrow_fdw_export_cupy</code>の返却するGPUバッファの識別子(text)を引数として受け取り、これをPython側では<code>cupy.ndarray</code>にマッピングして参照します。シェル上でスクリプトを実行する場合と大きな違いはありません。</p>
<p>ただ一点だけ、GPUカーネルを実行して<code>X</code>の列ごとの総和を<code>Y</code>に格納した後、<code>X</code>に0を代入して明示的に<code>cupy.ndarray</code>を解放するようにしています。
これは、PL/Pythonでスクリプトを実行した場合に<code>X</code>に格納した<code>cupy.ndarray</code>オブジェクトが生き続けてしまい、識別子で参照したGPUバッファがマップされ続けてしまうためのワークアラウンドです。</p>
<p><strong>PL/Pythonユーザ定義関数の例</strong></p>
<pre><code>CREATE OR REPLACE FUNCTION custom_average(x_ident text)
RETURNS float[] AS
$$
import cupy
import cupy_strom

X = cupy_strom.ipc_import(x_ident)
nattrs = X.shape[0]
nitems = X.shape[1]
gridSz = (nitems + 2047) &gt;&gt; 11;

Y = cupy.zeros((nattrs))

source='''
extern &quot;C&quot; __global__
           __launch_bounds__(1024)
void
kern_gpu_sum(double *y, const float *x, int nitems)
{
    __shared__ float lvalues[2048];
    int     gridSz = (nitems + 2047) / 2048;
    int     colIdx = blockIdx.x / gridSz;
    int     rowBase = (blockIdx.x % gridSz) * 2048;
    int     localId = 2 * threadIdx.x;
    int     i, k;

    // Load values to local shared buffer
    x += colIdx * nitems;
    for (i=threadIdx.x; i &lt; 2048; i+=blockDim.x)
        lvalues[i] = (rowBase + i &lt; nitems ? x[rowBase + i] : 0.0);
    __syncthreads();

    // Run reduction operations
    for (k=0; k &lt; 11; k++)
    {
        int     mask = ((1 &lt;&lt; k) - 1);

        if ((threadIdx.x &amp; mask) == 0)
            lvalues[localId] += lvalues[localId + (1&lt;&lt;k)];
        __syncthreads();
    }
    // Write back the total sum
    if (threadIdx.x == 0)
        atomicAdd(&amp;y[colIdx], lvalues[0]);
}
'''
kern = cupy.RawKernel(source, 'kern_gpu_sum')
kern.__call__((gridSz * nattrs,0,0),
              (1024,0,0),
              (Y,X,nitems))
X = 0   # unmap GPU memory

return Y / nitems
$$ LANGUAGE 'plpython3u';
</code></pre>

<p>PL/Pythonでユーザ定義関数を実装し、カスタムのGPUカーネルを実行する場合、Pythonスクリプトを呼び出すための一連のステップが異なってきます。</p>
<p>シェル上でスクリプトを実行した時とは異なり、SQLの実行中にユーザ定義関数を通じてPythonスクリプトを実行できるため、わざわざセッションを張ったり、SELECT文を実行して識別子を取得する必要はありません。</p>
<pre><code>=# SELECT custom_average(pgstrom.arrow_fdw_export_cupy('ft','{x,y,z}'::text[]));
                    custom_average
-------------------------------------------------------
 {50.0225953554688,-24.9964806318359,100.037490859375}
(1 row)
</code></pre>

<p>上記の例では、<code>pgstrom.arrow_fdw_export_cupy</code>の実行結果であるGPUバッファの識別子が、直接ユーザ定義関数<code>custom_average</code>の引数として入力され、そこでPythonスクリプトとの間でデータ交換が行われています。その後、GPUカーネルが起動され、結果を呼び出し元に返すまでの流れは全く同一です。</p>
<p>外部テーブルftからデータを読み出し、これを引数としてPL/Pythonユーザ定義関数に渡すのとは、PostgreSQL側で取り扱うべきデータサイズが全く異なってくる事に留意してください。
GPUバッファを介したデータ交換メカニズムの場合、これはある種の"ポインタ渡し"として機能するため、"値渡し"スタイルと比べてユーザ定義関数の呼び出しそのものが非常に軽い処理になっています。</p>
              
            </div>
          </div>
          <footer>
  
    <div class="rst-footer-buttons" role="navigation" aria-label="footer navigation">
      
        <a href="../ref_types/" class="btn btn-neutral float-right" title="データ型">Next <span class="icon icon-circle-arrow-right"></span></a>
      
      
        <a href="../gstore_fdw/" class="btn btn-neutral" title="GPUメモリストア"><span class="icon icon-circle-arrow-left"></span> Previous</a>
      
    </div>
  

  <hr/>

  <div role="contentinfo">
    <!-- Copyright etc -->
    
  </div>

  Built with <a href="http://www.mkdocs.org">MkDocs</a> using a <a href="https://github.com/snide/sphinx_rtd_theme">theme</a> provided by <a href="https://readthedocs.org">Read the Docs</a>.
</footer>
      
        </div>
      </div>

    </section>

  </div>

  <div class="rst-versions" role="note" style="cursor: pointer">
    <span class="rst-current-version" data-toggle="rst-current-version">
      
      
        <span><a href="../gstore_fdw/" style="color: #fcfcfc;">&laquo; Previous</a></span>
      
      
        <span style="margin-left: 15px"><a href="../ref_types/" style="color: #fcfcfc">Next &raquo;</a></span>
      
    </span>
</div>
    <script>var base_url = '..';</script>
    <script src="../js/theme.js" defer></script>
      <script src="../search/main.js" defer></script>

</body>
</html>
